//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : matv.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef MATV_H
#define	MATV_H

#include "kernels/matv_kernels.h"

namespace rrbsolver
{

template <class T>
inline void matv_phase1(T* dY, T* dX, const Grid orig)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, dX, &desc, 
        orig.nx, orig.ny, orig.nx * sizeof(T)), 
            "rrbsolver::bind texture @ kernel_matv1");
#endif

    dimGrid.x = orig.cx / NUM_THREADS_X;
    dimGrid.y = orig.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_matv1<T> <<< dimGrid, dimBlock >>> (dY, dX, orig);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_matv1");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_matv1");
#endif
}

template <class T>
inline void matv_phase2(T* dY, T* dX, const Grid orig)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texRef, dY, &desc, 
        orig.nx, orig.ny, orig.nx * sizeof(T)), 
            "rrbsolver::bind texture @ kernel_matv2");    
#endif

    dimGrid.x = orig.cx / NUM_THREADS_X;
    dimGrid.y = orig.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_matv2<T> <<< dimGrid, dimBlock >>> (dY, dX, orig); 
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_matv2");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texRef), 
        "rrbsolver::unbind texture @ kernel_matv2");
#endif
}

} // namespace rrbsolver

#endif // MATV_H

